{
 "cells": [
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "<div align=\"center\"><h1>使用 CUDA C/C++ 加速应用程序</h1></div>"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "![CUDA](./images/CUDA_Logo.jpg)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "加速计算正在取代 CPU 计算，成为最佳计算做法。加速计算带来的层出不穷的突破性进展、对加速应用程序日益增长的需求、轻松编写加速计算的编程规范以及支持加速计算的硬件的不断改进，所有这一切都在推动计算方式必然会过渡到加速计算。\n",
    "\n",
    "无论是从出色的性能还是易用性来看，[CUDA](https://developer.nvidia.com/about-cuda) 计算平台均是加速计算的制胜法宝。CUDA 提供一种可扩展 C、C++、Python 和 Fortran 等语言的编码范式，能够在世界上性能超强劲的并行处理器 NVIDIA GPU 上运行大量经加速的并行代码。CUDA 可以毫不费力地大幅加速应用程序，具有适用于 [DNN](https://developer.nvidia.com/cudnn)、[BLAS](https://developer.nvidia.com/cublas)、[图形分析](https://developer.nvidia.com/nvgraph) 和 [FFT](https://developer.nvidia.com/cufft) 等的高度优化库生态系统，并且还附带功能强大的 [命令行](http://docs.nvidia.com/cuda/profiler-users-guide/index.html#nvprof-overview) 和 [可视化分析器] (http://docs.nvidia.com/cuda/profiler-users-guide/index.html#visual)。\n",
    "\n",
    "CUDA 支持以下领域的许多（即便不是大多数）[世界上性能超强劲的应用程序](https://www.nvidia.com/en-us/data-center/gpu-accelerated-applications/catalog/?product_category_id=58,59,60,293,98,172,223,227,228,265,487,488,114,389,220,258,461&search=)：[计算流体动力学](https://www.nvidia.com/en-us/data-center/gpu-accelerated-applications/catalog/?product_category_id=10,12,16,17,19,51,53,71,87,121,124,156,157,195,202,203,204,312,339,340,395,407,448,485,517,528,529,541,245,216,104,462,513,250,492,420,429,490,10,12,16,17,19,51,53,71,87,121,124,156,157,195,202,203,204,312,339,340,395,407,448,485,517,528,529,541,245,216,104,462,513,250,492,420,429,490,10,12,16,17,19,51,53,71,87,121,124,156,157,195,202,203,204,312,339,340,395,407,448,485,517,528,529,541,245,216,104,462,513,250,492,420,429,490&search=)、[分子动力学](https://www.nvidia.com/en-us/data-center/gpu-accelerated-applications/catalog/?product_category_id=8,57,92,123,211,213,237,272,274,282,283,307,325,337,344,345,351,362,365,380,396,398,400,435,507,508,519,8,57,92,123,211,213,237,272,274,282,283,307,325,337,344,345,351,362,365,380,396,398,400,435,507,508,519,8,57,92,123,211,213,237,272,274,282,283,307,325,337,344,345,351,362,365,380,396,398,400,435,507,508,519,8,57,92,123,211,213,237,272,274,282,283,307,325,337,344,345,351,362,365,380,396,398,400,435,507,508,519&search=)、[量子化学](https://www.nvidia.com/en-us/data-center/gpu-accelerated-applications/catalog/?product_category_id=8,57,92,123,211,213,237,272,274,282,283,307,325,337,344,345,351,362,365,380,396,398,400,435,507,508,519,8,57,92,123,211,213,237,272,274,282,283,307,325,337,344,345,351,362,365,380,396,398,400,435,507,508,519&search=)、[物理学](https://www.nvidia.com/en-us/data-center/gpu-accelerated-applications/catalog/?product_category_id=6,24,116,118,119,135,229,231,372,373,392,393,489,493,494,495,496,497,498,67,170,216,281,6,24,116,118,119,135,229,231,372,373,392,393,489,493,494,495,496,497,498,67,170,216,281,6,24,116,118,119,135,229,231,372,373,392,393,489,493,494,495,496,497,498,67,170,216,281,6,24,116,118,119,135,229,231,372,373,392,393,489,493,494,495,496,497,498,67,170,216,281,6,24,116,118,119,135,229,231,372,373,392,393,489,493,494,495,496,497,498,67,170,216,281&search=) 和高性能计算 (HPC)。\n",
    "\n",
    "学习 CUDA 将能助您加速自己的应用程序。加速应用程序的执行速度远远超过 CPU 应用程序，并且可以执行 CPU 应用程序受限于其性能而无法执行的计算。在本实验中, 您将学习使用 CUDA C/C++ 为加速应用程序编程的入门知识，这些入门知识足以让您开始加速自己的 CPU 应用程序以获得性能提升并助您迈入全新的计算领域。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "## 如要充分利用本实验,您应已能胜任如下任务：\n",
    "\n",
    "To get the most out of this lab you should already be able to:\n",
    "\n",
    "- 在 C 中声明变量、编写循环并使用 if/else 语句。\n",
    "- 在 C 中定义和调用函数。\n",
    "- 在 C 中分配数组。\n",
    "\n",
    "无需 CUDA 预备知识。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "## Objectives\n",
    "\n",
    "当您在本实验完成学习后，您将能够：\n",
    "\n",
    "- 编写、编译及运行既可调用 CPU 函数也可**启动** GPU **核函数*的 C/C++ 程序。\n",
    "- 使用**执行配置**控制并行**线程层次结构**。\n",
    "- 重构串行循环以在 GPU 上并行执行其迭代。\n",
    "- 分配和释放可用于 CPU 和 GPU 的内存。\n",
    "- 处理 CUDA 代码生成的错误。\n",
    "- 加速 CPU 应用程序。\n"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "## Accelerated Systems\n",
    "\n",
    "*加速系统*又称*异构系统*，由 CPU 和 GPU 组成。加速系统会运行 CPU 程序，这些程序也会转而启动将受益于 GPU 大规模并行计算能力的函数。本实验环境是一个包含 NVIDIA GPU 的加速系统。可以使用 `nvidia-smi` (*Systems Management Interface*) 命令行命令查询有关此 GPU 的信息。现在，可以在下方的代码执行单元上使用 `CTRL` + `ENTER` 发出 `nvidia-smi` 命令。无论您何时需要执行代码，均可在整个实验中找到这些单元。代码运行后，运行该命令的输出将打印在代码执行单元的正下方。在运行下方的代码执行块后，请注意在输出中找到并记录 GPU 的名称。"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "!nvidia-smi"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "## GPU-accelerated Vs. CPU-only Applications\n",
    "\n",
    "以下幻灯片将直观呈现即将发布的材料的概要信息。点击浏览一遍这些幻灯片，然后再继续深入了解以下章节中的主题。\n",
    "\n",
    "<script>console.log('hi');</script>"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "%%HTML\n",
    "\n",
    "<div align=\"center\"><iframe src=\"https://view.officeapps.live.com/op/view.aspx?src=https://developer.download.nvidia.com/training/courses/C-AC-01-V1/AC_CUDA_C-zh/AC_CUDA_C_1-zh.pptx\" frameborder=\"0\" width=\"900\" height=\"550\" allowfullscreen=\"true\" mozallowfullscreen=\"true\" webkitallowfullscreen=\"true\"></iframe></div>"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "## Writing Application Code for the GPU\n",
    "\n",
    "CUDA 为许多常用编程语言提供扩展，而在本实验中，我们将会为 C/C++ 提供扩展。这些语言扩展可让开发人员在 GPU 上轻松运行其源代码中的函数。\n",
    "\n",
    "以下是一个 `.cu` 文件（`.cu` 是 CUDA 加速程序的文件扩展名）。其中包含两个函数，第一个函数将在 CPU 上运行，第二个将在 GPU 上运行。请抽点时间找出这两个函数在定义方式和调用方式上的差异。\n",
    "\n",
    "```cpp\n",
    "void CPUFunction()\n",
    "{\n",
    "  printf(\"This function is defined to run on the CPU.\\n\");\n",
    "}\n",
    "\n",
    "__global__ void GPUFunction()\n",
    "{\n",
    "  printf(\"This function is defined to run on the GPU.\\n\");\n",
    "}\n",
    "\n",
    "int main()\n",
    "{\n",
    "  CPUFunction();\n",
    "\n",
    "  GPUFunction<<<1, 1>>>();\n",
    "  cudaDeviceSynchronize();\n",
    "}\n",
    "```\n",
    "\n",
    "以下是一些需要特别注意的重要代码行，以及加速计算中使用的一些其他常用术语：\n",
    "\n",
    "`__global__ void GPUFunction()`\n",
    "  - `__global__` 关键字表明以下函数将在 GPU 上运行并可**全局**调用，而在此种情况下，则指由 CPU 或 GPU 调用。\n",
    "  - 通常，我们将在 CPU 上执行的代码称为**主机**代码，而将在 GPU 上运行的代码称为**设备**代码。\n",
    "  - 注意返回类型为 `void`。使用 `__global__` 关键字定义的函数需要返回 `void` 类型。\n",
    "\n",
    "`GPUFunction<<<1, 1>>>();`\n",
    "  - 通常，当调用要在 GPU 上运行的函数时，我们将此种函数称为**已启动**的**核函数**。\n",
    "  - 启动核函数时，我们必须提供**执行配置**，即在向核函数传递任何预期参数之前使用 `<<< ... >>>` 语法完成的配置。\n",
    "  - 在宏观层面，程序员可通过执行配置为核函数启动指定**线程层次结构**，从而定义线程组（称为**线程块**）的数量，以及要在每个线程块中执行的**线程**数量。稍后将在本实验深入探讨执行配置，但现在请注意正在使用包含 `1` 线程（第二个配置参数）的 `1` 线程块（第一个执行配置参数）启动核函数。\n",
    "\n",
    "`cudaDeviceSynchronize();`\n",
    "  - 与许多 C/C++ 代码不同，核函数启动方式为**异步**：CPU 代码将继续执行*而无需等待核函数完成启动*。\n",
    "  - 调用 CUDA 运行时提供的函数 `cudaDeviceSynchronize` 将导致主机 (CPU) 代码暂作等待，直至设备 (GPU) 代码执行完成，才能在 CPU 上恢复执行。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "### Exercise: Write a Hello GPU Kernel\n",
    "\n",
    "[`01-hello-gpu.cu`](../../../../../edit/tasks/task1/task/01_AC_CUDA_C-zh/01-hello/01-hello-gpu.cu)（*<---- 点击源文件链接，以在另一个标签中打开并进行编辑*）包含已在运行的程序。其中包含两个函数，都有打印 \\”Hello from the CPU\\” 消息。您的目标是重构源文件中的 `helloGPU` 函数，以便该函数实际上在 GPU 上运行，并打印指示执行此操作的消息。\n",
    "\n",
    "- 请先重构应用程序，然后使用下方的 `nvcc` 命令编译和运行该应用程序（请记住，您可以通过使用 `CTRL + ENTER` 启动该应用程序以执行代码执行单元的内容）。[`01-hello-gpu.cu`](../../../../../edit/tasks/task1/task/01_AC_CUDA_C-zh/01-hello/01-hello-gpu.cu) 中的注释将有助您完成操作。如您遇到问题或要检查自己的操作，请参阅 [解决方案](../../../../../edit/tasks/task1/task/01_AC_CUDA_C-zh/01-hello/solutions/01-hello-gpu-solution.cu)。在使用下方命令进行编译和运行之前，请不要忘记保存对文件所作的更改。"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "!nvcc -arch=sm_70 -o hello-gpu 01-hello/01-hello-gpu.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "成功重构 [`01-hello-gpu.cu`](../../../../../edit/tasks/task1/task/01_AC_CUDA_C-zh/01-hello/01-hello-gpu.cu) 后，请进行以下修改，并尝试在每次更改后编译并运行该应用程序（通过使用 `CTRL + ENTER` 点击上方的代码执行单元）。若出现错误，请花时间仔细阅读错误内容：熟知错误内容会在您开始编写自己的加速代码时，给与很大的帮助。\n",
    "\n",
    "- 从核函数定义中删除关键字 `__global__`。注意错误中的行号：您认为错误中的 \\”configured\\” 是什么意思？完成后，请替换 `__global__`。\n",
    "- 移除执行配置：您对 \\”configured\\” 的理解是否仍旧合理？完成后，请替换执行配置。\n",
    "- 移除对 `cudaDeviceSynchronize` 的调用。在编译和运行代码之前，猜猜会发生什么情况，可以回顾一下核函数采取的是异步启动，且 `cudaDeviceSynchronize` 会使主机执行暂作等待，直至核函数执行完成后才会继续。完成后，请替换对 `cudaDeviceSynchronize` 的调用。\n",
    "- 重构 `01-hello-gpu.cu`，以便 `Hello from the GPU` 在 `Hello from the CPU` **之前**打印。\n",
    "- 重构 `01-hello-gpu.cu`，以便 `Hello from the GPU` 打印**两次**，一次是在 `Hello from the CPU` **之前**，另一次是在 `Hello from the CPU` **之后**。\n"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "### Compiling and Running Accelerated CUDA Code\n",
    "\n",
    "本节包含上述为编译和运行 `.cu` 程序而调用的`nvcc` 命令的详细信息。\n",
    "\n",
    "CUDA 平台附带 [**NVIDIA CUDA 编译器**](http://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html) `nvcc`，可以编译 CUDA 加速应用程序，其中包含主机和设备代码。就本实验而言，`nvcc` 的讨论范围将根据我们的迫切需求据实确定。完成本实验学习后，有意深究 `nvcc` 的所有用户均可从 [文档](http://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html) 开始入手。\n",
    "\n",
    "曾使用过 `gcc` 的用户会对 `nvcc` 感到非常熟悉。例如，编译 `some-CUDA.cu` 文件就很简单：\n",
    "\n",
    "`nvcc -arch=sm_70 -o out some-CUDA.cu -run`\n",
    "  - `nvcc` 是使用 `nvcc` 编译器的命令行命令。\n",
    "  - 将 `some-CUDA.cu` 作为文件传递以进行编译。\n",
    "  - `o` 标志用于指定编译程序的输出文件。\n",
    "  - `arch` 标志表示该文件必须编译为哪个**架构**类型。本示例中，`sm_70` 将用于专门针对本实验运行的 Volta GPU 进行编译，但有意深究的用户可以参阅有关 [`arch` 标志](http://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html#options-for-steering-gpu-code-generation)、[虚拟架构特性] (http://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html#gpu-feature-list) 和 [GPU特性](http://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html#gpu-feature-list) 的文档。\n",
    "\n",
    "  - 为方便起见，提供 `run` 标志将执行已成功编译的二进制文件。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "## CUDA Thread Hierarchy\n",
    "\n",
    "以下幻灯片将直观呈现即将发布的材料的概要信息。点击浏览一遍这些幻灯片，然后再继续深入了解以下章节中的主题。"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "%%HTML\n",
    "\n",
    "<div align=\"center\"><iframe src=\"https://view.officeapps.live.com/op/view.aspx?src=https://developer.download.nvidia.com/training/courses/C-AC-01-V1/AC_CUDA_C-zh/AC_CUDA_C_2-zh.pptx\" frameborder=\"0\" width=\"900\" height=\"550\" allowfullscreen=\"true\" mozallowfullscreen=\"true\" webkitallowfullscreen=\"true\"></iframe></div>"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "## Launching Parallel Kernels\n",
    "\n",
    "程序员可通过执行配置指定有关如何启动核函数以在多个 GPU **线程**中并行运行的详细信息。更准确地说，程序员可通过执行配置指定线程组（称为**线程块**或简称为**块**）数量以及其希望每个线程块所包含的线程数量。执行配置的语法如下：\n",
    "\n",
    "`<<< NUMBER_OF_BLOCKS, NUMBER_OF_THREADS_PER_BLOCK>>>`\n",
    "\n",
    "**启动核函数时，核函数代码由每个已配置的线程块中的每个线程执行**。\n",
    "\n",
    "因此，如果假设已定义一个名为 `someKernel` 的核函数，则下列情况为真：\n",
    "  - `someKernel<<<1, 1>>()` 配置为在具有单线程的单个线程块中运行后，将只运行一次。\n",
    "  - `someKernel<<<1, 10>>()` 配置为在具有 10 线程的单个线程块中运行后，将运行 10 次。\n",
    "  - `someKernel<<<10, 1>>()` 配置为在 10 个线程块（每个均具有单线程）中运行后，将运行 10 次。\n",
    "  - `someKernel<<<10, 10>>()` 配置为在 10 个线程块（每个均具有 10 线程）中运行后，将运行 100 次。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "### Exercise: Launch Parallel Kernels\n",
    "\n",
    "[`01-first-parallel.cu`](../../../../../edit/tasks/task1/task/01_AC_CUDA_C-zh/02-first-parallel/01-basic-parallel.cu) 目前已作出十分基本的函数调用，即打印消息 `This should be running in parallel`。请按下列步骤首先进行重构使之在 GPU 上运行，然后在单个线程块中并行运行，最后则在多个线程块中运行。如您遇到问题，请参阅 [解决方案](../../../../../edit/tasks/task1/task/01_AC_CUDA_C-zh/02-first-parallel/solutions/01-basic-parallel-solution.cu)。\n",
    "\n",
    "- 重构 `firstParallel` 函数以便在 GPU 上作为 CUDA 核函数启动。在使用下方 `nvcc` 命令编译和运行 `01-basic-parallel.cu` 后，您应仍能看到函数的输出。\n",
    "- 重构 `firstParallel` 核函数以便在 5 个线程中并行执行，且均在同一个线程块中执行。在编译和运行代码后，您应能看到输出消息已打印 5 次。\n",
    "- 再次重构 `firstParallel` 核函数，并使其在 5 个线程块内并行执行（每个线程块均包含 5 个线程）。编译和运行之后，您应能看到输出消息现已打印 25 次。"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "!nvcc -arch=sm_70 -o basic-parallel 02-first-parallel/01-basic-parallel.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "\n",
    "## CUDA-Provided Thread Hierarchy Variables\n",
    "\n",
    "以下幻灯片将直观呈现即将发布的材料的概要信息。点击浏览一遍这些幻灯片，然后再继续深入了解以下章节中的主题。"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "%%HTML\n",
    "\n",
    "<div align=\"center\"><iframe src=\"https://view.officeapps.live.com/op/view.aspx?src=https://developer.download.nvidia.com/training/courses/C-AC-01-V1/AC_CUDA_C-zh/AC_CUDA_C_3-zh.pptx\" frameborder=\"0\" width=\"900\" height=\"550\" allowfullscreen=\"true\" mozallowfullscreen=\"true\" webkitallowfullscreen=\"true\"></iframe></div>"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "## Thread and Block Indices\n",
    "\n",
    "每个线程在其线程块内部均会被分配一个索引，从 `0` 开始。此外，每个线程块也会被分配一个索引，并从 `0` 开始。正如线程组成线程块，线程块又会组成**网格**，而网格是 CUDA 线程层次结构中级别最高的实体。简言之，CUDA 核函数在由一个或多个线程块组成的网格中执行，且每个线程块中均包含相同数量的一个或多个线程。\n",
    "\n",
    "CUDA 核函数可以访问能够识别如下两种索引的特殊变量：正在执行核函数的线程（位于线程块内）索引和线程所在的线程块（位于网格内）索引。这两种变量分别为 `threadIdx.x` 和 `blockIdx.x`。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "### Exercise: Use Specific Thread and Block Indices\n",
    "\n",
    "目前，[`01-thread-and-block-idx.cu`](../../../../../edit/tasks/task1/task/01_AC_CUDA_C-zh/03-indices/01-thread-and-block-idx.cu) 文件包含一个正在打印失败消息的执行中的核函数。打开文件以了解如何更新执行配置，以便打印成功消息。重构后，请使用下方代码执行单元编译并运行代码以确认您的工作。如您遇到问题，请参阅 [解决方案](../../../../../edit/tasks/task1/task/01_AC_CUDA_C-zh/03-indices/solutions/01-thread-and-block-idx-solution.cu)。"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "!nvcc -arch=sm_70 -o thread-and-block-idx 03-indices/01-thread-and-block-idx.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "## Accelerating For Loops\n",
    "\n",
    "对 CPU 应用程序中的循环进行加速的时机已经成熟：我们并非要顺次运行循环的每次迭代，而是让每次迭代都在自身线程中并行运行。考虑以下“for 循环”，尽管很明显，但还是请注意它控制着循环将执行的次数，并会界定循环的每次迭代将会发生的情况：\n",
    "\n",
    "```cpp\n",
    "int N = 2<<20;\n",
    "for (int i = 0; i < N; ++i)\n",
    "{\n",
    "  printf(\"%d\\n\", i);\n",
    "}\n",
    "```\n",
    "\n",
    "如要并行此循环，必须执行以下 2 个步骤：\n",
    "\n",
    "- 必须编写完成**循环的单次迭代**工作的核函数。\n",
    "- 由于核函数与其他正在运行的核函数无关，因此执行配置必须使核函数执行正确的次数，例如循环迭代的次数。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "### Exercise: Accelerating a For Loop with a Single Block of Threads\n",
    "\n",
    "目前，[`01-single-block-loop.cu`](../../../../../edit/tasks/task1/task/01_AC_CUDA_C-zh/04-loops/01-single-block-loop.cu) 内的 `loop` 函数运行着一个“for 循环”并将连续打印 `0` 至 `9` 之间的所有数字。将 `loop` 函数重构为 CUDA 核函数，使其在启动后并行执行 `N` 次迭代。重构成功后，应仍能打印 `0` 至 `9` 之间的所有数字。如您遇到问题，请参阅 [解决方案](../../../../../edit/tasks/task1/task/01_AC_CUDA_C-zh/04-loops/solutions/01-single-block-loop-solution.cu)。"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "!nvcc -arch=sm_70 -o single-block-loop 04-loops/01-single-block-loop.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "## Coordinating Parallel Threads\n",
    "\n",
    "以下幻灯片将直观呈现即将发布的材料的概要信息。点击浏览一遍这些幻灯片，然后再继续深入了解以下章节中的主题。"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "%%HTML\n",
    "\n",
    "<div align=\"center\"><iframe src=\"https://view.officeapps.live.com/op/view.aspx?src=https://developer.download.nvidia.com/training/courses/C-AC-01-V1/AC_CUDA_C-zh/AC_CUDA_C_4-zh.pptx\" frameborder=\"0\" width=\"900\" height=\"550\" allowfullscreen=\"true\" mozallowfullscreen=\"true\" webkitallowfullscreen=\"true\"></iframe></div>"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "## Using Block Dimensions for More Parallelization\n",
    "\n",
    "线程块包含的线程具有数量限制：确切地说是 1024 个。为增加加速应用程序中的并行量，我们必须要能在多个线程块之间进行协调。\n",
    "\n",
    "CUDA 核函数可以访问给出块中线程数的特殊变量：`blockDim.x`。通过将此变量与 `blockIdx.x` 和 `threadIdx.x` 变量结合使用，并借助惯用表达式 `threadIdx.x + blockIdx.x * blockDim.x` 在包含多个线程的多个线程块之间组织并行执行，并行性将得以提升。以下是详细示例。\n",
    "\n",
    "执行配置 `<<<10, 10>>>` 将启动共计拥有 100 个线程的网格，这些线程均包含在由 10 个线程组成的 10 个线程块中。因此，我们希望每个线程（`0` 至 `99` 之间）都能计算该线程的某个唯一索引。\n",
    "\n",
    "- 如果线程块 `blockIdx.x` 等于 `0`，则 `blockIdx.x * blockDim.x` 为 `0`。向 `0` 添加可能的 `threadIdx.x` 值（`0` 至 `9`），之后便可在包含 100 个线程的网格内生成索引 `0` 至 `9`。\n",
    "- 如果线程块 `blockIdx.x` 等于 `1`，则 `blockIdx.x * blockDim.x` 为 `10`。向 `10` 添加可能的 `threadIdx.x` 值（`0` 至 `9`），之后便可在包含 100 个线程的网格内生成索引 `10` 至 `19`。\n",
    "- 如果线程块 `blockIdx.x` 等于 `5`，则 `blockIdx.x * blockDim.x` 为 `50`。向 `50` 添加可能的 `threadIdx.x` 值（`0` 至 `9`），之后便可在包含 100 个线程的网格内生成索引 `50` 至 `59`。\n",
    "- 如果线程块 `blockIdx.x` 等于 `9`，则 `blockIdx.x * blockDim.x` 为 `90`。向 `90` 添加可能的 `threadIdx.x` 值（`0` 至 `9`），之后便可在包含 100 个线程的网格内生成索引 `90` 至 `99`。\n"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "### Exercise: Accelerating a For Loop with Multiple Blocks of Threads\n",
    "\n",
    "目前，[`02-multi-block-loop.cu`](../../../../../edit/tasks/task1/task/01_AC_CUDA_C-zh/04-loops/02-multi-block-loop.cu) 内的 `loop` 函数运行着一个“for 循环”并将连续打印 `0` 至 `9` 之间的所有数字。将 `loop` 函数重构为 CUDA 核函数，使其在启动后并行执行 `N` 次迭代。重构成功后，应仍能打印 `0` 至 `9` 之间的所有数字。对于本练习，作为附加限制，请使用启动*至少 2 个线程块*的执行配置。如您遇到问题，请参阅 [解决方案](../../../../../edit/tasks/task1/task/01_AC_CUDA_C-zh/04-loops/solutions/02-multi-block-loop-solution.cu)。"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "!nvcc -arch=sm_70 -o multi-block-loop 04-loops/02-multi-block-loop.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "## Allocating Memory to be accessed on the GPU and the CPU\n",
    "\n",
    "CUDA 的最新版本（版本 6 和更高版本）已能轻松分配可用于 CPU 主机和任意数量 GPU 设备的内存。尽管现今有许多适用于内存管理并可支持加速应用程序中最优性能的 [中高级技术](http://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html#memory-optimizations)，但我们现在要介绍的基础 CUDA 内存管理技术不但能够支持远超 CPU 应用程序的卓越性能，而且几乎不会产生任何开发人员成本。\n",
    "\n",
    "如要分配和释放内存，并获取可在主机和设备代码中引用的指针，请使用 `cudaMallocManaged` 和 `cudaFree` 取代对 `malloc` 和 `free` 的调用，如下例所示：\n",
    "\n",
    "```cpp\n",
    "// CPU-only\n",
    "\n",
    "int N = 2<<20;\n",
    "size_t size = N * sizeof(int);\n",
    "\n",
    "int *a;\n",
    "a = (int *)malloc(size);\n",
    "\n",
    "// Use `a` in CPU-only program.\n",
    "\n",
    "free(a);\n",
    "```\n",
    "\n",
    "```cpp\n",
    "// Accelerated\n",
    "\n",
    "int N = 2<<20;\n",
    "size_t size = N * sizeof(int);\n",
    "\n",
    "int *a;\n",
    "// Note the address of `a` is passed as first argument.\n",
    "cudaMallocManaged(&a, size);\n",
    "\n",
    "// Use `a` on the CPU and/or on any GPU in the accelerated system.\n",
    "\n",
    "cudaFree(a);\n",
    "```"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "### Exercise: Array Manipulation on both the Host and Device\n",
    "\n",
    "[`01-double-elements.cu`](../../../../../edit/tasks/task1/task/01_AC_CUDA_C-zh/05-allocate/01-double-elements.cu) 程序会分配一个数组、在主机上使用整数值对其进行初始化并尝试在 GPU 上对其中的每个值并行加倍，然后在主机上确认加倍操作是否成功。目前，程序将无法执行：因其正尝试在主机和设备上与指针 `a` 指向的数组进行交互，但仅分配可在主机上访问的数组（使用 `malloc`）。重构应用程序以满足以下条件，如您遇到问题，请参阅 [解决方案](../../../../../edit/tasks/task1/task/01_AC_CUDA_C-zh/05-allocate/solutions/01-double-elements-solution.cu)：\n",
    "\n",
    "- 指针 `a` 应可供主机和设备代码使用。\n",
    "- 应该正确释放指针 `a` 的内存。"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "!nvcc -arch=sm_70 -o double-elements 05-allocate/01-double-elements.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "## Grid Size Work Amount Mismatch\n",
    "\n",
    "以下幻灯片将直观呈现即将发布的材料的概要信息。点击浏览一遍这些幻灯片，然后再继续深入了解以下章节中的主题。"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "%%HTML\n",
    "\n",
    "<div align=\"center\"><iframe src=\"https://view.officeapps.live.com/op/view.aspx?src=https://developer.download.nvidia.com/training/courses/C-AC-01-V1/AC_CUDA_C-zh/AC_CUDA_C_5-zh.pptx\" frameborder=\"0\" width=\"900\" height=\"550\" allowfullscreen=\"true\" mozallowfullscreen=\"true\" webkitallowfullscreen=\"true\"></iframe></div>"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "## Handling Block Configuration Mismatches to Number of Needed Threads\n",
    "\n",
    "可能会出现这样的情况，即无法表示会创建并行循环所需确切线程数量的执行配置。\n",
    "\n",
    "常见示例与希望选择的最佳线程块大小有关。例如，鉴于 GPU 的硬件特性，所含线程的数量为 32 的倍数的线程块是为理想的选择，因其具备性能上的优势。假设我们要启动一些线程块且每个线程块中均包含 256 个线程（32 的倍数），并需运行 1000 个并行任务（此处使用极小的数量以便于说明），则任何数量的线程块均无法在网格中精确生成 1000 个总线程，因为没有任何整数值在乘以 32 后可以恰好等于 1000。\n",
    "\n",
    "可以通过以下方式轻松解决这种情况：\n",
    "\n",
    "- 编写执行配置，使其创建的线程数**超过**执行分配工作所需的线程数。\n",
    "- 将值作为参数传递到核函数 (`N`) 中，以表示要处理的数据集总大小或完成工作所需的总线程数。\n",
    "- 计算网格内的线程索引后（使用 `tid+bid*bdim`），请检查该索引是否超过 `N`，并且只在不超过的情况下执行与核函数相关的工作。\n",
    "\n",
    "\n",
    "以下是编写执行配置的惯用方法示例，适用于 `N` 和线程块中的线程数已知，但无法保证网格中的线程数和 `N` 之间完全匹配的情况。如此一来，便可确保网格中至少始终拥有 `N` 所需的线程数，且超出的线程数至多仅可相当于 1 个线程块的线程数量：\n",
    "\n",
    "```cpp\n",
    "// Assume `N` is known\n",
    "int N = 100000;\n",
    "\n",
    "// Assume we have a desire to set `threads_per_block` exactly to `256`\n",
    "size_t threads_per_block = 256;\n",
    "\n",
    "// Ensure there are at least `N` threads in the grid, but only 1 block's worth extra\n",
    "size_t number_of_blocks = (N + threads_per_block - 1) / threads_per_block;\n",
    "\n",
    "some_kernel<<<number_of_blocks, threads_per_block>>>(N);\n",
    "```\n",
    "\n",
    "由于上述执行配置致使网格中的线程数超过 `N`，因此需要注意 `some_kernel` 定义中的内容，以确保 `some_kernel` 在由其中一个 \\”extra\\” 线程执行时不会尝试访问超出范围的数据元素：\n",
    "\n",
    "```cpp\n",
    "__global__ some_kernel(int N)\n",
    "{\n",
    "  int idx = threadIdx.x + blockIdx.x * blockDim.x;\n",
    "\n",
    "  if (idx < N) // Check to make sure `idx` maps to some value within `N`\n",
    "  {\n",
    "    // Only do work if it does\n",
    "  }\n",
    "}\n",
    "```"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "### Exercise: Accelerating a For Loop with a Mismatched Execution Configuration\n",
    "\n",
    "[`02-mismatched-config-loop.cu`](../../../../../edit/tasks/task1/task/01_AC_CUDA_C-zh/05-allocate/02-mismatched-config-loop.cu) 中的程序使用 `cudaMallocManaged` 为包含 1000 个元素的整数数组分配内存，然后试图使用 CUDA 核函数以并行方式初始化数组中的所有值。此程序假设 `N` 和 `threads_per_block` 的数量均为已知。您的任务是完成以下两个目标，如您遇到问题，请参阅 [解决方案](../../../../../edit/tasks/task1/task/01_AC_CUDA_C-zh/05-allocate/solutions/02-mismatched-config-loop-solution.cu)：\n",
    "\n",
    "- 为 `number_of_blocks` 分配一个值，以确保线程数至少与指针 `a` 中可供访问的元素数同样多。\n",
    "- 更新 `initializeElementsTo` 核函数以确保不会尝试访问超出范围的数据元素。"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "!nvcc -arch=sm_70 -o mismatched-config-loop 05-allocate/02-mismatched-config-loop.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "## Grid-Stride Loops\n",
    "\n",
    "以下幻灯片将直观呈现即将发布的材料的概要信息。点击浏览一遍这些幻灯片，然后再继续深入了解以下章节中的主题。"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "%%HTML\n",
    "\n",
    "<div align=\"center\"><iframe src=\"https://view.officeapps.live.com/op/view.aspx?src=https://developer.download.nvidia.com/training/courses/C-AC-01-V1/AC_CUDA_C-zh/AC_CUDA_C_1-zh.pptx\" frameborder=\"0\" width=\"900\" height=\"550\" allowfullscreen=\"true\" mozallowfullscreen=\"true\" webkitallowfullscreen=\"true\"></iframe></div>"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "## Data Sets Larger then the Grid\n",
    "\n",
    "或出于选择，为了要创建具有超高性能的执行配置，或出于需要，一个网格中的线程数量可能会小于数据集的大小。请思考一下包含 1000 个元素的数组和包含 250 个线程的网格（此处使用极小的规模以便于说明）。此网格中的每个线程将需使用 4 次。如要实现此操作，一种常用方法便是在核函数中使用**网格跨度循环**。\n",
    "\n",
    "在网格跨度循环中，每个线程将在网格内使用 `tid+bid*bdim` 计算自身唯一的索引，并对数组内该索引的元素执行相应运算，然后将网格中的线程数添加到索引并重复此操作，直至超出数组范围。例如，对于包含 500 个元素的数组和包含 250 个线程的网格，网格中索引为 20 的线程将执行如下操作：\n",
    "\n",
    "- 对包含 500 个元素的数组的元素 20 执行相应运算\n",
    "- 将其索引增加 250，使网格的大小达到 270\n",
    "- 对包含 500 个元素的数组的元素 270 执行相应运算\n",
    "- 将其索引增加 250，使网格的大小达到 520\n",
    "- 由于 520 现已超出数组范围，因此线程将停止工作\n",
    "\n",
    "\n",
    "CUDA 提供一个可给出网格中线程块数的特殊变量：`gridDim.x`。然后计算网格中的总线程数，即网格中的线程块数乘以每个线程块中的线程数：`gridDim.x * blockDim.x`。带着这样的想法来看看以下核函数中网格跨度循环的详细示例：\n",
    "\n",
    "```cpp\n",
    "__global void kernel(int *a, int N)\n",
    "{\n",
    "  int indexWithinTheGrid = threadIdx.x + blockIdx.x * blockDim.x;\n",
    "  int gridStride = gridDim.x * blockDim.x;\n",
    "\n",
    "  for (int i = indexWithinTheGrid; i < N; i += gridStride)\n",
    "  {\n",
    "    // do work on a[i];\n",
    "  }\n",
    "}\n",
    "```"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "### Exercise: Use a Grid-Stride Loop to Manipulate an Array Larger than the Grid\n",
    "\n",
    "重构 [`03-grid-stride-double.cu`](../../../../../edit/tasks/task1/task/01_AC_CUDA_C-zh/05-allocate/03-grid-stride-double.cu) 以在 `doubleElements` 核函数中使用网格跨度循环，进而使小于 `N` 的网格可以重用线程以覆盖数组中的每个元素。程序会打印数组中的每个元素是否均已加倍，而当前该程序会准确打印出 `FALSE`。如您遇到问题，请参阅 [解决方案](../../../../../edit/tasks/task1/task/01_AC_CUDA_C-zh/05-allocate/solutions/03-grid-stride-double-solution.cu)。"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "!nvcc -arch=sm_70 -o grid-stride-double 05-allocate/03-grid-stride-double.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "## Error Handling\n",
    "\n",
    "与在任何应用程序中一样，加速 CUDA 代码中的错误处理同样至关重要。即便不是大多数，也有许多 CUDA 函数（例如，[内存管理函数](http://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__MEMORY.html#group__CUDART__MEMORY)）会返回类型为 `cudaError_t` 的值，该值可用于检查调用函数时是否发生错误。以下是对调用 `cudaMallocManaged` 函数执行错误处理的示例：\n",
    "\n",
    "```cpp\n",
    "cudaError_t err;\n",
    "err = cudaMallocManaged(&a, N)                    // Assume the existence of `a` and `N`.\n",
    "\n",
    "if (err != cudaSuccess)                           // `cudaSuccess` is provided by CUDA.\n",
    "{\n",
    "  printf(\"Error: %s\\n\", cudaGetErrorString(err)); // `cudaGetErrorString` is provided by CUDA.\n",
    "}\n",
    "```\n",
    "\n",
    "启动定义为返回 `void` 的核函数后，将不会返回类型为 `cudaError_t` 的值。为检查启动核函数时是否发生错误（例如，如果启动配置错误），CUDA 提供 `cudaGetLastError` 函数，该函数会返回类型为 `cudaError_t` 的值。\n",
    "\n",
    "```cpp\n",
    "/*\n",
    " * This launch should cause an error, but the kernel itself\n",
    " * cannot return it.\n",
    " */\n",
    "\n",
    "someKernel<<<1, -1>>>();  // -1 is not a valid number of threads.\n",
    "\n",
    "cudaError_t err;\n",
    "err = cudaGetLastError(); // `cudaGetLastError` will return the error from above.\n",
    "if (err != cudaSuccess)\n",
    "{\n",
    "  printf(\"Error: %s\\n\", cudaGetErrorString(err));\n",
    "}\n",
    "```\n",
    "\n",
    "最后，为捕捉异步错误（例如，在异步核函数执行期间），请务必检查后续同步 CUDA 运行时 API 调用所返回的状态（例如 `cudaDeviceSynchronize`）；如果之前启动的其中一个核函数失败，则将返回错误。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "### Exercise: Add Error Handling\n",
    "\n",
    "目前，[`01-add-error-handling.cu`](../../../../../edit/tasks/task1/task/01_AC_CUDA_C-zh/06-errors/01-add-error-handling.cu) 会编译、运行并打印已加倍失败的数组元素。不过，该程序不会指明其中是否存在任何错误。重构应用程序以处理 CUDA 错误，以便您可以了解程序出现的问题并进行有效调试。您将需要调查在调用 CUDA 函数时可能出现的同步错误，以及在执行 CUDA 核函数时可能出现的异步错误。如您遇到问题，请参阅 [解决方案](../../../../../edit/tasks/task1/task/01_AC_CUDA_C-zh/06-errors/solutions/01-add-error-handling-solution.cu)。"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "!nvcc -arch=sm_70 -o add-error-handling 06-errors/01-add-error-handling.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "### CUDA Error Handling Function\n",
    "\n",
    "创建一个包装 CUDA 函数调用的宏对于检查错误十分有用。以下是一个宏示例，您可以在余下练习中随时使用：\n",
    "\n",
    "```cpp\n",
    "#include <stdio.h>\n",
    "#include <assert.h>\n",
    "\n",
    "inline cudaError_t checkCuda(cudaError_t result)\n",
    "{\n",
    "  if (result != cudaSuccess) {\n",
    "    fprintf(stderr, \"CUDA Runtime Error: %s\\n\", cudaGetErrorString(result));\n",
    "    assert(result == cudaSuccess);\n",
    "  }\n",
    "  return result;\n",
    "}\n",
    "\n",
    "int main()\n",
    "{\n",
    "\n",
    "/*\n",
    " * The macro can be wrapped around any function returning\n",
    " * a value of type `cudaError_t`.\n",
    " */\n",
    "\n",
    "  checkCuda( cudaDeviceSynchronize() )\n",
    "}\n",
    "```"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "## Summary\n",
    "\n",
    "至此，您已经完成以下列出的所有实验学习目标：\n",
    "\n",
    "- 编写、编译及运行既可调用 CPU 函数也可**启动** GPU **核函数*的 C/C++ 程序。\n",
    "- 使用**执行配置**控制并行**线程层次结构**。\n",
    "- 重构串行循环以在 GPU 上并行执行其迭代。\n",
    "- 分配和释放可用于 CPU 和 GPU 的内存。\n",
    "- 处理 CUDA 代码生成的错误。\n",
    "\n",
    "\n",
    "现在，您将完成实验的最终目标：\n",
    "\n",
    "- 加速 CPU 应用程序。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "### Final Exercise: Accelerate Vector Addition Application\n",
    "\n",
    "下面的挑战将使您有机会运用在实验中学到的知识。其中涉及加速 CPU 向量加法程序，尽管该程序不甚复杂，但仍能让您有机会重点运用所学的借助 CUDA 加速 GPU 应用程序的相关知识。完成此练习后，如果您有富余时间并有意深究，可继续学习*高阶内容*部分以了解涉及更复杂代码库的一些挑战。\n",
    "\n",
    "[`01-vector-add.cu`](../../../../../edit/tasks/task1/task/01_AC_CUDA_C-zh/07-vector-add/01-vector-add.cu) 包含一个可正常运作的 CPU 向量加法应用程序。加速其 `addVectorsInto` 函数，使之在 GPU 上以 CUDA 核函数运行并使其并行执行工作。鉴于需发生以下操作，如您遇到问题，请参阅 [解决方案](../../../../../edit/tasks/task1/task/01_AC_CUDA_C-zh/07-vector-add/solutions/01-vector-add-solution.cu)。\n",
    "\n",
    "- 扩充 `addVectorsInto` 定义，使之成为 CUDA 核函数。\n",
    "- 选择并使用有效的执行配置，以使 `addVectorsInto` 作为 CUDA 核函数启动。\n",
    "- 更新内存分配，内存释放以反映主机和设备代码需要访问 3 个向量：`a`、`b` 和 `result`。\n",
    "- 重构 `addVectorsInto` 的主体：将在单个线程内部启动，并且只需对输入向量执行单线程操作。确保线程从不尝试访问输入向量范围之外的元素，并注意线程是否需对输入向量的多个元素执行操作。\n",
    "- 在 CUDA 代码可能以其他方式静默失败的位置添加错误处理。"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "!nvcc -arch=sm_70 -o vector-add 07-vector-add/01-vector-add.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "## Advanced Content\n",
    "\n",
    "以下练习为时间富余且有意深究的学习者提供额外挑战。这些挑战需要使用更先进的技术加以应对，并且其提供的有用知识也很少。因此，完成这些挑战着实不易，但您在此过程中亦会收获重大进步。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "## Grids and Blocks of 2 and 3 Dimensions\n",
    "\n",
    "可以将网格和线程块定义为最多具有 3 个维度。使用多个维度定义网格和线程块绝不会对其性能造成任何影响，但这在处理具有多个维度的数据时可能非常有用，例如 2D 矩阵。如要定义二维或三维网格或线程块，可以使用 CUDA 的 `dim3` 类型，即如下所示：\n",
    "\n",
    "```cpp\n",
    "dim3 threads_per_block(16, 16, 1);\n",
    "dim3 number_of_blocks(16, 16, 1);\n",
    "someKernel<<<number_of_blocks, threads_per_block>>>();\n",
    "```\n",
    "\n",
    "鉴于以上示例，`someKernel` 内部的变量 `gridDim.x`、`gridDim.y`、`blockDim.x` 和 `blockDim.y` 均将等于 `16`。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "### Exercise: Accelerate 2D Matrix Multiply Application\n",
    "\n",
    "文件 [`01-matrix-multiply-2d.cu`](../../../../../edit/tasks/task1/task/01_AC_CUDA_C-zh/08-matrix-multiply/01-matrix-multiply-2d.cu) 包含一个功能齐全的主机函数 `matrixMulCPU`。您的任务是扩建 CUDA 核函数 `matrixMulGPU`。源代码将使用这两个函数执行矩阵乘法，并比较它们的答案以验证您编写的 CUDA 核函数是否正确。使用以下指南获得操作支持，如您遇到问题，请参阅 [解决方案](../../../../../edit/tasks/task1/task/01_AC_CUDA_C-zh/08-matrix-multiply/solutions/01-matrix-multiply-2d-solution.cu)：\n",
    "\n",
    "- 您将需创建执行配置，其参数均为 `dim3` 值，且 `x` 和 `y` 维度均设为大于 `1`。\n",
    "- 在核函数主体内部，您将需要按照惯例在网格内建立所运行线程的唯一索引，但应为线程建立两个索引：一个用于网格的 x 轴，另一个用于网格的 y 轴。"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "!nvcc -arch=sm_70 -o matrix-multiply-2d 08-matrix-multiply/01-matrix-multiply-2d.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "### Exercise: Accelerate A Thermal Conductivity Application\n",
    "\n",
    "在下面的练习中，您将为模拟金属银二维热传导的应用程序执行加速操作。\n",
    "\n",
    "将 [`01-heat-conduction.cu`](../../../../../edit/tasks/task1/task/01_AC_CUDA_C-zh/09-heat/01-heat-conduction.cu) 内的 `step_kernel_mod` 函数转换为在 GPU 上执行，并修改 `main` 函数以恰当分配在 CPU 和 GPU 上使用的数据。`step_kernel_ref` 函数在 CPU 上执行并用于检查错误。由于此代码涉及浮点计算，因此不同的处理器甚或同一处理器上的简单重排操作都可能导致结果略有出入。为此，错误检查代码会使用错误阈值，而非查找完全匹配。如您遇到问题，请参阅 [解决方案](../../../../../edit/tasks/task1/task/01_AC_CUDA_C-zh/09-heat/solutions/01-heat-conduction-solution.cu)。"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "!nvcc -arch=sm_70 -o heat-conduction 09-heat/01-heat-conduction.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "> 此任务中的原始热传导 CPU 源代码取自于休斯顿大学的文章 [An OpenACC Example Code for a C-based heat conduction code](http://docplayer.net/30411068-An-openacc-example-code-for-a-c-based-heat-conduction-code.html)（基于 C 的热传导代码的 OpenACC 示例代码）。"
   ]
  }
 ],
 "metadata": {
  "kernelspec": {
   "display_name": "Python 2",
   "language": "python",
   "name": "python2"
  },
  "language_info": {
   "codemirror_mode": {
    "name": "ipython",
    "version": 2
   },
   "file_extension": ".py",
   "mimetype": "text/x-python",
   "name": "python",
   "nbconvert_exporter": "python",
   "pygments_lexer": "ipython2",
   "version": "2.7.12"
  }
 },
 "nbformat": 4,
 "nbformat_minor": 2
}
